{
 "cells": [
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "默认与非默认流"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "#include <stdio.h>\n",
    "#include <unistd.h>\n",
    "\n",
    "__global__ void printNumber(int number)\n",
    "{\n",
    "  printf(\"%d\\n\", number);\n",
    "}\n",
    "\n",
    "int main()\n",
    "{\n",
    "  for (int i = 0; i < 5; ++i)\n",
    "  {\n",
    "    cudaStream_t stream;\n",
    "    cudaStreamCreate(&stream);\n",
    "    printNumber<<<1, 1, 0, stream>>>(i);\n",
    "    cudaStreamDestroy(stream);\n",
    "  }\n",
    "  cudaDeviceSynchronize();\n",
    "}\n",
    "\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Exercise: Overlap Kernel Execution and Memory Copy Back to Host"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "#include <stdio.h>\n",
    "\n",
    "__global__\n",
    "void initWith(float num, float *a, int N)\n",
    "{\n",
    "\n",
    "  int index = threadIdx.x + blockIdx.x * blockDim.x;\n",
    "  int stride = blockDim.x * gridDim.x;\n",
    "\n",
    "  for(int i = index; i < N; i += stride)\n",
    "  {\n",
    "    a[i] = num;\n",
    "  }\n",
    "}\n",
    "\n",
    "__global__\n",
    "void addVectorsInto(float *result, float *a, float *b, int N)\n",
    "{\n",
    "  int index = threadIdx.x + blockIdx.x * blockDim.x;\n",
    "  int stride = blockDim.x * gridDim.x;\n",
    "\n",
    "  for(int i = index; i < N; i += stride)\n",
    "  {\n",
    "    result[i] = a[i] + b[i];\n",
    "  }\n",
    "}\n",
    "\n",
    "void checkElementsAre(float target, float *vector, int N)\n",
    "{\n",
    "  for(int i = 0; i < N; i++)\n",
    "  {\n",
    "    if(vector[i] != target)\n",
    "    {\n",
    "      printf(\"FAIL: vector[%d] - %0.0f does not equal %0.0f\\n\", i, vector[i], target);\n",
    "      exit(1);\n",
    "    }\n",
    "  }\n",
    "  printf(\"Success! All values calculated correctly.\\n\");\n",
    "}\n",
    "\n",
    "int main()\n",
    "{\n",
    "  int deviceId;\n",
    "  int numberOfSMs;\n",
    "\n",
    "  cudaGetDevice(&deviceId);\n",
    "  cudaDeviceGetAttribute(&numberOfSMs, cudaDevAttrMultiProcessorCount, deviceId);\n",
    "\n",
    "  const int N = 2<<24;\n",
    "  size_t size = N * sizeof(float);\n",
    "\n",
    "  float *a;\n",
    "  float *b;\n",
    "  float *c;\n",
    "  float *h_c;\n",
    "\n",
    "  cudaMalloc(&a, size);\n",
    "  cudaMalloc(&b, size);\n",
    "  cudaMalloc(&c, size);\n",
    "  cudaMallocHost(&h_c, size);\n",
    "\n",
    "  size_t threadsPerBlock;\n",
    "  size_t numberOfBlocks;\n",
    "\n",
    "  threadsPerBlock = 256;\n",
    "  numberOfBlocks = 32 * numberOfSMs;\n",
    "\n",
    "  cudaError_t addVectorsErr;\n",
    "  cudaError_t asyncErr;\n",
    "\n",
    "  /*\n",
    "   * Create 3 streams to run initialize the 3 data vectors in parallel.\n",
    "   */\n",
    "\n",
    "  cudaStream_t stream1, stream2, stream3;\n",
    "  cudaStreamCreate(&stream1);\n",
    "  cudaStreamCreate(&stream2);\n",
    "  cudaStreamCreate(&stream3);\n",
    "\n",
    "  /*\n",
    "   * Give each `initWith` launch its own non-standard stream.\n",
    "   */\n",
    "\n",
    "  initWith<<<numberOfBlocks, threadsPerBlock, 0, stream1>>>(3, a, N);\n",
    "  initWith<<<numberOfBlocks, threadsPerBlock, 0, stream2>>>(4, b, N);\n",
    "  initWith<<<numberOfBlocks, threadsPerBlock, 0, stream3>>>(0, c, N);\n",
    "  cudaDeviceSynchronize();\n",
    "\n",
    "  for (int i = 0; i < 4; ++i)\n",
    "  {\n",
    "    cudaStream_t stream;\n",
    "    cudaStreamCreate(&stream);\n",
    "\n",
    "    addVectorsInto<<<numberOfBlocks/4, threadsPerBlock, 0, stream>>>(&c[i*N/4], &a[i*N/4], &b[i*N/4], N/4);\n",
    "    cudaMemcpyAsync(&h_c[i*N/4], &c[i*N/4], size/4, cudaMemcpyDeviceToHost, stream);\n",
    "    cudaStreamDestroy(stream);\n",
    "  }\n",
    "\n",
    "  cudaMemcpy(h_c, c, size, cudaMemcpyDeviceToHost);\n",
    "\n",
    "  addVectorsErr = cudaGetLastError();\n",
    "  if(addVectorsErr != cudaSuccess) printf(\"Error: %s\\n\", cudaGetErrorString(addVectorsErr));\n",
    "\n",
    "  asyncErr = cudaDeviceSynchronize();\n",
    "  if(asyncErr != cudaSuccess) printf(\"Error: %s\\n\", cudaGetErrorString(asyncErr));\n",
    "\n",
    "  checkElementsAre(7, h_c, N);\n",
    "\n",
    "  /*\n",
    "   * Destroy streams when they are no longer needed.\n",
    "   */\n",
    "\n",
    "  cudaStreamDestroy(stream1);\n",
    "  cudaStreamDestroy(stream2);\n",
    "  cudaStreamDestroy(stream3);\n",
    "\n",
    "  cudaFree(a);\n",
    "  cudaFree(b);\n",
    "  cudaFree(c);\n",
    "  cudaFreeHost(h_c);\n",
    "}\n"
   ]
  }
 ],
 "metadata": {
  "language_info": {
   "name": "python"
  },
  "orig_nbformat": 4
 },
 "nbformat": 4,
 "nbformat_minor": 2
}
